#include <stdio.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

using namespace std;

static __global__ void testKernel(float* array) {
    array[threadIdx.x] = threadIdx.x;
}

int main() {
    int num = 5;
    float* array = nullptr;
    cudaMallocHost(&array, sizeof(float) * num);  // 申请host侧的Pinned Memory
    // Pinned Memory正确使用场景是适合大数据量传输的中间存储，即从host侧拷贝到device侧。比如一个大权重文件，先放在Pinned Memory中，然后再拷贝到device侧。
    testKernel<<<1, num>>>(array);  // 此处传给核函数的是host侧内存，核函数在device侧执行时可以直接读写host侧内存，但是效率较低，不能频繁读写。
    cudaDeviceSynchronize();  // 等待device侧所有线程都执行完
    for (int i = 0; i < num; i++) {
        printf("array[%d]=%f\n", i, array[i]);
    }

    cudaFreeHost(array);  // 使用cudaMallocHost分配的内存，用cudaFreeHost释放

    return 0;
}

// array[0]=0.000000
// array[1]=1.000000
// array[2]=2.000000
// array[3]=3.000000
// array[4]=4.000000